<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub::BlockExchange&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt; Class Template Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li class="current"><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="annotated.html"><span>Class&#160;List</span></a></li>
      <li><a href="classes.html"><span>Class&#160;Index</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div id="nav-path" class="navpath">
  <ul>
<li class="navelem"><a class="el" href="namespacecub.html">cub</a></li><li class="navelem"><a class="el" href="classcub_1_1_block_exchange.html">BlockExchange</a></li>  </ul>
</div>
</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="classcub_1_1_block_exchange-members.html">List of all members</a>  </div>
  <div class="headertitle">
<div class="title">cub::BlockExchange&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt; Class Template Reference<div class="ingroups"><a class="el" href="group___block_module.html">Block-wide</a></div></div>  </div>
</div><!--header-->
<div class="contents">
<a name="details" id="details"></a><h2 class="groupheader">Detailed description</h2>
<div class="textblock"><h3>template&lt;
    typename T, 
    int BLOCK_THREADS, 
    int ITEMS_PER_THREAD, 
    bool WARP_TIME_SLICING = false&gt;<br/>
class cub::BlockExchange&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;</h3>

<p>The <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">BlockExchange</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for rearranging data partitioned across a CUDA thread block. </p>
<div class="image">
<img src="transpose_logo.png" alt="transpose_logo.png"/>
<div class="caption">
.</div></div>
 <dl class="section user"><dt>Overview</dt><dd>It is commonplace for blocks of threads to rearrange data items between threads. For example, the global memory subsystem prefers access patterns where data items are "striped" across threads (where consecutive threads access consecutive items), yet most block-wide operations prefer a "blocked" partitioning of items across threads (where consecutive items belong to a single thread).</dd></dl>
<dl class="section user"><dt></dt><dd><a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">BlockExchange</a> supports the following types of data exchanges:<ul>
<li>Transposing between <a href="index.html#sec5sec4"><em>blocked</em></a> and <a href="index.html#sec5sec4"><em>striped</em></a> arrangements</li>
<li>Transposing between <a href="index.html#sec5sec4"><em>blocked</em></a> and <a href="index.html#sec5sec4"><em>warp-striped</em></a> arrangements</li>
<li>Scattering ranked items to a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a></li>
<li>Scattering ranked items to a <a href="index.html#sec5sec4"><em>striped arrangement</em></a></li>
</ul>
</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">T</td><td>The data type to be exchanged. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td>The number of items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">WARP_TIME_SLICING</td><td><b>[optional]</b> When <code>true</code>, only use enough shared memory for a single warp's worth of tile data, time-slicing the block-wide exchange over multiple synchronized rounds. Yields a smaller memory footprint at the expense of decreased parallelism. (Default: false)</td></tr>
  </table>
  </dd>
</dl>
<dl class="section user"><dt>A Simple Example</dt><dd>Every thread in the block uses the <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">BlockExchange</a> class by first specializing the <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">BlockExchange</a> type, then instantiating an instance with parameters for communication, and finally invoking collective member functions. </dd></dl>
<dl class="section user"><dt></dt><dd>The code snippet below illustrates the conversion from a "blocked" to a "striped" arrangement of 512 integer items partitioned across 128 threads where each thread owns 4 items. </dd></dl>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockExchange for 128 threads owning 4 integer items each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockExchange</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockExchange::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Load a tile of data striped across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    cub::LoadStriped&lt;LOAD_DEFAULT, 128&gt;(threadIdx.x, d_data, thread_data);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively exchange data into a blocked arrangement across threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>(temp_storage).StripedToBlocked(thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of striped input <code>thread_data</code> across the block of threads is <code>{ [0,128,256,384], [1,129,257,385], ..., [127,255,383,511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>.</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Proper device-specific padding ensures zero bank conflicts for most types. </li>
</ul>
</dd></dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00110">110</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>
</div><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html">TempStorage</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The operations exposed by <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">BlockExchange</a> require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the <code>__shared__</code> keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or <code>union</code>'d with other storage allocation types to facilitate memory reuse.  <a href="structcub_1_1_block_exchange_1_1_temp_storage.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="pub-methods"></a>
Public Methods</h2></td></tr>
<tr><td colspan="2"><div class="groupHeader">Collective constructors</div></td></tr>
<tr class="memitem:a31f5a889713220dfe21669a4865f3f4e"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="a31f5a889713220dfe21669a4865f3f4e"></a>
__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e">BlockExchange</a> ()</td></tr>
<tr class="memdesc:a31f5a889713220dfe21669a4865f3f4e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage. Threads are identified using <code>threadIdx.x</code>. <br/></td></tr>
<tr class="separator:a31f5a889713220dfe21669a4865f3f4e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aff7137dd257eeec8d7628b4f9f4e1457"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#aff7137dd257eeec8d7628b4f9f4e1457">BlockExchange</a> (<a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage)</td></tr>
<tr class="memdesc:aff7137dd257eeec8d7628b4f9f4e1457"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>.  <a href="#aff7137dd257eeec8d7628b4f9f4e1457">More...</a><br/></td></tr>
<tr class="separator:aff7137dd257eeec8d7628b4f9f4e1457"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:aa0689007bc51632404a1893e52a0e6e9"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#aa0689007bc51632404a1893e52a0e6e9">BlockExchange</a> (int linear_tid)</td></tr>
<tr class="memdesc:aa0689007bc51632404a1893e52a0e6e9"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#aa0689007bc51632404a1893e52a0e6e9">More...</a><br/></td></tr>
<tr class="separator:aa0689007bc51632404a1893e52a0e6e9"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a166150541730a842f1af039e34131ea9"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a166150541730a842f1af039e34131ea9">BlockExchange</a> (<a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html">TempStorage</a> &amp;temp_storage, int linear_tid)</td></tr>
<tr class="memdesc:a166150541730a842f1af039e34131ea9"><td class="mdescLeft">&#160;</td><td class="mdescRight">Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier.  <a href="#a166150541730a842f1af039e34131ea9">More...</a><br/></td></tr>
<tr class="separator:a166150541730a842f1af039e34131ea9"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Structured exchanges</div></td></tr>
<tr class="memitem:a063505380e68bb5744d12fa7e1f16145"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a063505380e68bb5744d12fa7e1f16145">StripedToBlocked</a> (T items[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:a063505380e68bb5744d12fa7e1f16145"><td class="mdescLeft">&#160;</td><td class="mdescRight">Transposes data items from <em>striped</em> arrangement to <em>blocked</em> arrangement.  <a href="#a063505380e68bb5744d12fa7e1f16145">More...</a><br/></td></tr>
<tr class="separator:a063505380e68bb5744d12fa7e1f16145"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a5e487b2f594fa0ae466bf1c7b6b9c06d"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a5e487b2f594fa0ae466bf1c7b6b9c06d">BlockedToStriped</a> (T items[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:a5e487b2f594fa0ae466bf1c7b6b9c06d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Transposes data items from <em>blocked</em> arrangement to <em>striped</em> arrangement.  <a href="#a5e487b2f594fa0ae466bf1c7b6b9c06d">More...</a><br/></td></tr>
<tr class="separator:a5e487b2f594fa0ae466bf1c7b6b9c06d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a09969c12793c1876a2e6328c97c0330a"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a09969c12793c1876a2e6328c97c0330a">WarpStripedToBlocked</a> (T items[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:a09969c12793c1876a2e6328c97c0330a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Transposes data items from <em>warp-striped</em> arrangement to <em>blocked</em> arrangement.  <a href="#a09969c12793c1876a2e6328c97c0330a">More...</a><br/></td></tr>
<tr class="separator:a09969c12793c1876a2e6328c97c0330a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a74fc1a62cac21f54289bb7092250a868"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a74fc1a62cac21f54289bb7092250a868">BlockedToWarpStriped</a> (T items[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:a74fc1a62cac21f54289bb7092250a868"><td class="mdescLeft">&#160;</td><td class="mdescRight">Transposes data items from <em>blocked</em> arrangement to <em>warp-striped</em> arrangement.  <a href="#a74fc1a62cac21f54289bb7092250a868">More...</a><br/></td></tr>
<tr class="separator:a74fc1a62cac21f54289bb7092250a868"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Scatter exchanges</div></td></tr>
<tr class="memitem:ac8dd22c90a28e54a8c1d6b0637054661"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#ac8dd22c90a28e54a8c1d6b0637054661">ScatterToBlocked</a> (T items[ITEMS_PER_THREAD], int ranks[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ac8dd22c90a28e54a8c1d6b0637054661"><td class="mdescLeft">&#160;</td><td class="mdescRight">Exchanges data items annotated by rank into <em>blocked</em> arrangement.  <a href="#ac8dd22c90a28e54a8c1d6b0637054661">More...</a><br/></td></tr>
<tr class="separator:ac8dd22c90a28e54a8c1d6b0637054661"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a53516886ba6750060819581b1bb23d68"><td class="memItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html#a53516886ba6750060819581b1bb23d68">ScatterToStriped</a> (T items[ITEMS_PER_THREAD], int ranks[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:a53516886ba6750060819581b1bb23d68"><td class="mdescLeft">&#160;</td><td class="mdescRight">Exchanges data items annotated by rank into <em>striped</em> arrangement.  <a href="#a53516886ba6750060819581b1bb23d68">More...</a><br/></td></tr>
<tr class="separator:a53516886ba6750060819581b1bb23d68"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Constructor &amp; Destructor Documentation</h2>
<a class="anchor" id="aff7137dd257eeec8d7628b4f9f4e1457"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::<a class="el" href="classcub_1_1_block_exchange.html">BlockExchange</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage. Threads are identified using <code>threadIdx.x</code>. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html" title="The operations exposed by BlockExchange require a temporary memory allocation of this nested type for...">TempStorage</a> </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00660">660</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="aa0689007bc51632404a1893e52a0e6e9"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::<a class="el" href="classcub_1_1_block_exchange.html">BlockExchange</a> </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using a private static allocation of shared memory as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00674">674</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a166150541730a842f1af039e34131ea9"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::<a class="el" href="classcub_1_1_block_exchange.html">BlockExchange</a> </td>
          <td>(</td>
          <td class="paramtype"><a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html">TempStorage</a> &amp;&#160;</td>
          <td class="paramname"><em>temp_storage</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Collective constructor using the specified memory allocation as temporary storage. Each thread is identified using the supplied linear thread identifier. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">temp_storage</td><td>Reference to memory allocation having layout type <a class="el" href="structcub_1_1_block_exchange_1_1_temp_storage.html" title="The operations exposed by BlockExchange require a temporary memory allocation of this nested type for...">TempStorage</a> </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td><b>[optional]</b> A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00688">688</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Member Function Documentation</h2>
<a class="anchor" id="a5e487b2f594fa0ae466bf1c7b6b9c06d"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::BlockedToStriped </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD])</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Transposes data items from <em>blocked</em> arrangement to <em>striped</em> arrangement. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the conversion from a "blocked" to a "striped" arrangement of 512 integer items partitioned across 128 threads where each thread owns 4 items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockExchange for 128 threads owning 4 integer items each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockExchange</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockExchange::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively exchange data into a striped arrangement across threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>(temp_storage).BlockedToStriped(thread_data);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Store data striped across block threads into an ordered tile</span></div>
<div class="line">    cub::StoreStriped&lt;STORE_DEFAULT, 128&gt;(threadIdx.x, d_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of blocked input <code>thread_data</code> across the block of threads is <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,128,256,384], [1,129,257,385], ..., [127,255,383,511] }</code> in preparation for storing to global memory. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">items</td><td>Items to exchange, converting between <em>blocked</em> and <em>striped</em> arrangements. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00784">784</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a74fc1a62cac21f54289bb7092250a868"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::BlockedToWarpStriped </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD])</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Transposes data items from <em>blocked</em> arrangement to <em>warp-striped</em> arrangement. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the conversion from a "blocked" to a "warp-striped" arrangement of 512 integer items partitioned across 128 threads where each thread owns 4 items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockExchange for 128 threads owning 4 integer items each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockExchange</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockExchange::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Obtain a segment of consecutive items that are blocked across threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    ...</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively exchange data into a warp-striped arrangement across threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>(temp_storage).BlockedToWarpStriped(thread_data);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Store data striped across warp threads into an ordered tile</span></div>
<div class="line">    cub::StoreStriped&lt;STORE_DEFAULT, 128&gt;(threadIdx.x, d_data, thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of blocked input <code>thread_data</code> across the block of threads is <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,32,64,96], [1,33,65,97], [2,34,66,98], ..., [415,447,479,511] }</code> in preparation for storing to global memory. (The first 128 items are striped across the first warp of 32 threads, the second 128 items are striped across the second warp, etc.) </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">items</td><td>Items to exchange, converting between <em>blocked</em> and <em>warp-striped</em> arrangements. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00872">872</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ac8dd22c90a28e54a8c1d6b0637054661"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::ScatterToBlocked </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>ranks</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Exchanges data items annotated by rank into <em>blocked</em> arrangement. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">items</td><td>Items to exchange </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">ranks</td><td>Corresponding scatter ranks </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00891">891</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a53516886ba6750060819581b1bb23d68"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::ScatterToStriped </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>ranks</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Exchanges data items annotated by rank into <em>striped</em> arrangement. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed. </p>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">items</td><td>Items to exchange </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">ranks</td><td>Corresponding scatter ranks </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00904">904</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a063505380e68bb5744d12fa7e1f16145"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::StripedToBlocked </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD])</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Transposes data items from <em>striped</em> arrangement to <em>blocked</em> arrangement. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the conversion from a "striped" to a "blocked" arrangement of 512 integer items partitioned across 128 threads where each thread owns 4 items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockExchange for 128 threads owning 4 integer items each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockExchange</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockExchange::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Load a tile of ordered data into a striped arrangement across block threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    cub::LoadStriped&lt;LOAD_DEFAULT, 128&gt;(threadIdx.x, d_data, thread_data);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively exchange data into a blocked arrangement across threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>(temp_storage).StripedToBlocked(thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of striped input <code>thread_data</code> across the block of threads is <code>{ [0,128,256,384], [1,129,257,385], ..., [127,255,383,511] }</code> after loading from global memory. The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">items</td><td>Items to exchange, converting between <em>striped</em> and <em>blocked</em> arrangements. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00740">740</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a09969c12793c1876a2e6328c97c0330a"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;typename T , int BLOCK_THREADS, int ITEMS_PER_THREAD, bool WARP_TIME_SLICING = false&gt; </div>
<table class="mlabels">
  <tr>
  <td class="mlabels-left">
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void <a class="el" href="classcub_1_1_block_exchange.html">cub::BlockExchange</a>&lt; T, BLOCK_THREADS, ITEMS_PER_THREAD, WARP_TIME_SLICING &gt;::WarpStripedToBlocked </td>
          <td>(</td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD])</td><td></td>
          <td></td>
        </tr>
      </table>
  </td>
  <td class="mlabels-right">
<span class="mlabels"><span class="mlabel">inline</span></span>  </td>
  </tr>
</table>
</div><div class="memdoc">

<p>Transposes data items from <em>warp-striped</em> arrangement to <em>blocked</em> arrangement. </p>
<p>A subsequent <code>__syncthreads()</code> threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., <code>temp_storage</code>) is to be reused or repurposed.</p>
<p>The code snippet below illustrates the conversion from a "warp-striped" to a "blocked" arrangement of 512 integer items partitioned across 128 threads where each thread owns 4 items. </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line">__global__ <span class="keywordtype">void</span> ExampleKernel(<span class="keywordtype">int</span> *d_data, ...)</div>
<div class="line">{</div>
<div class="line">    <span class="comment">// Specialize BlockExchange for 128 threads owning 4 integer items each</span></div>
<div class="line">    <span class="keyword">typedef</span> <a class="code" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange&lt;int, 128, 4&gt;</a> <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Allocate shared memory for BlockExchange</span></div>
<div class="line">    __shared__ <span class="keyword">typename</span> BlockExchange::TempStorage temp_storage;</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Load a tile of ordered data into a warp-striped arrangement across warp threads</span></div>
<div class="line">    <span class="keywordtype">int</span> thread_data[4];</div>
<div class="line">    cub::LoadSWarptriped&lt;LOAD_DEFAULT&gt;(threadIdx.x, d_data, thread_data);</div>
<div class="line"></div>
<div class="line">    <span class="comment">// Collectively exchange data into a blocked arrangement across threads</span></div>
<div class="line">    <a class="code" href="classcub_1_1_block_exchange.html#a31f5a889713220dfe21669a4865f3f4e" title="Collective constructor for 1D thread blocks using a private static allocation of shared memory as tem...">BlockExchange</a>(temp_storage).WarpStripedToBlocked(thread_data);</div>
</div><!-- fragment --> </dd></dl>
<dl class="section user"><dt></dt><dd>Suppose the set of warp-striped input <code>thread_data</code> across the block of threads is <code>{ [0,32,64,96], [1,33,65,97], [2,34,66,98], ..., [415,447,479,511] }</code> after loading from global memory. (The first 128 items are striped across the first warp of 32 threads, the second 128 items are striped across the second warp, etc.) The corresponding output <code>thread_data</code> in those threads will be <code>{ [0,1,2,3], [4,5,6,7], [8,9,10,11], ..., [508,509,510,511] }</code>. </dd></dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in,out]</td><td class="paramname">items</td><td>Items to exchange, converting between <em>warp-striped</em> and <em>blocked</em> arrangements. </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__exchange_8cuh_source.html#l00827">827</a> of file <a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a>.</p>

</div>
</div>
<hr/>The documentation for this class was generated from the following file:<ul>
<li><a class="el" href="block__exchange_8cuh_source.html">block_exchange.cuh</a></li>
</ul>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:11 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
